Summary

perf(attn): overlap the long-context KV load with compute in the GQA flash-decode split via cp.async double-buffering. The existing GQA split (fa_split_gqa_kernel) stages each KV tile into shared memory with a synchronous uint4 __ldg, hits a __syncthreads barrier, then reduces it — so the warps stall on the global KV read that dominates 4k/16k decode. This adds fa_split_gqa_pipe_kernel, which prefetches tile N+1 with cp.async into a second smem buffer while tile N is being reduced, hiding the KV-load latency behind the QK/softmax/V math.

What changed

• File: kernels/csrc/cuda/attention/flash_decode_split.cu (only file touched).
• New fa_split_gqa_pipe_kernel<HEAD_DIM, GQA, TILE> — a 2-stage cp.async pipeline over the KV tiles. 16-byte (uint4 = 8×bf16) .cg copies (same granularity as the current vectorized staging), issued via inline PTX so it stays header-free and NVRTC-device-only safe. A commit_group / wait_group pair keeps exactly one tile in flight; barriers are placed so the buffer being refilled is WAR-safe against the prior tile's reads.
• Selected by default for the GQA long-context path; SPARKINFER_FAPIPE=0 restores the single-buffer fa_split_gqa_kernel for A/B.
• Two K/V smem buffers double the per-tile shared memory to 4·TILE·128·2B (~14 KB) — well within the 48 KB default, no opt-in needed.

Correctness

The math is unchanged and the emitted partials are byte-identical to fa_split_gqa_kernel — the same KV values reach shared memory, only prefetched a tile earlier. So the top-1 / KL accuracy gate is unaffected by construction.

accuracy.sh (RTX 5090, Qwen3-30B-A3B Q4_K_M): top-1 = <0.xx>, KL = <0.0xx> (pass).

Proof of speedup

On-device eval runs only when the box is ticked and the table shows a real end-to-end decode gain (after > before, from bench/scripts/bench.sh).

• Tested on RTX 5090 (sm_120)

Env: RTX 5090 · sm_120 · CUDA <12.8/13> · Qwen3-30B-A3B Q4_K_M · 128 generated tokens.

Decode tok/s (end-to-end, from bench/scripts/bench.sh):

# paste bench/scripts/bench.sh output here (FAPIPE=0 baseline -> FAPIPE=1 this PR)